#include "common/book.h"
#include "common/gpu_anim.h"

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <cmath>
#include <string.h>
#include <stdlib.h>

#define DEMONO			2
#define DIMEXP			5
#define MASKLEN			(DIMEXP*2)
#define DIM				(1<<DIMEXP)
#define CANVDIM			256
#define CELLDIM			(CANVDIM/DIM)
#define GAP				1
#define R				2
#define	REPNO			9
#define REPNOPERLAYER	3
#define AGENTNO			(DIM*DIM)
#define LAYERNO			(REPNO / REPNOPERLAYER)
#define EMPTY			((1<<MASKLEN)-1)

#define TWO_TO_ONE(x,y) (x)+(y)*DIM 
#define POS_TO_EXIT_1D(pos, exit) abs(pos % DIM - exit % DIM) + abs(pos / DIM - exit / DIM)
#define POS_TO_EXIT_2D(x, y, exit) (x-exit%DIM) * (x-exit%DIM) + (y-exit/DIM) * (y-exit/DIM) 
//#define POS_TO_EXIT_2D(x, y, exit) abs(x - exit % DIM) + abs(y - exit % DIM)
inline __host__ void timer_start(CUevent &start, CUevent &stop){
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);
}
inline __host__ float timer_stop(CUevent &start, CUevent &stop){
	float elapsedTime = -1; 
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);
	return elapsedTime;
}
inline __device__ void encode(const int value, int &data, const int repid){
	int mask = (1<<MASKLEN)-1; 
	mask = mask<<(repid*MASKLEN);
	mask = ~mask;
	data = data & mask | value<<(repid*MASKLEN);
}
inline __device__ void decode(int &value, const int data, const int repid){
	int mask = (1<<MASKLEN*(repid+1))-1;
	value = data & mask; 
	value = value>>(MASKLEN*repid);
}
inline __host__ void decode_host(int &value, const int data, const int repid){
	int mask = (1<<MASKLEN*(repid+1))-1;
	value = data & mask; 
	value = value>>(MASKLEN*repid);
}

void generate_frame( uchar4 *pixels, void*, GPUAnimBitmap &bitmap );
void init_state(GPUAnimBitmap &bitmap);
void decision_point(int x, int y, int* nextPos, int *exitPos);
int	DEBUGREP;

int main( void ) {
	int a;
	printf("enter rep # for debug: ");
	scanf("%d", &a);
	DEBUGREP = a;
	remove("log.txt");
	GPUAnimBitmap  bitmap( DIM*CELLDIM, DIM*CELLDIM, LAYERNO * AGENTNO, NULL );
	printf("animate or not: ");
	scanf("%d", &a);
	bitmap.animated = a;
	init_state(bitmap);	
	bitmap.anim_and_exit(
		(void (*)(uchar4*,void*,GPUAnimBitmap &))generate_frame, 
		NULL,
		(void (*)(int,int,int*,int*))decision_point );
	printf("Success!\n");
#define host_data_type float
	host_data_type *hdata = (host_data_type*)malloc(REPNO*sizeof(host_data_type));
	cudaMemcpy(hdata, bitmap.runningTime, REPNO*sizeof(host_data_type), cudaMemcpyDeviceToHost);
	for(int i=0; i<REPNO; i++)
		printf("%f ", hdata[i]);
	return 0;
}

int		EXITNO		= REPNO;
__global__ void cleanup(bool *fini, int *loop, float *time, float oneIterTime){
	unsigned int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if(idx < REPNO){
		if(fini[idx] != true){
			loop[idx]++;
			time[idx] += oneIterTime;
		}
		fini[idx] = true;
	}
}

__global__ void wash(uchar4 *ptr){
	int x = threadIdx.x + blockIdx.x*blockDim.x;
	int y = threadIdx.y + blockIdx.y*blockDim.y;
	int offset = x + y * blockDim.x * gridDim.x;
	uchar4 color;
	color.x = 0;
	color.y = 0;
	color.z = 0;
	color.w = 255;
	if(offset<DIM*DIM*CELLDIM*CELLDIM)
		ptr[offset]=color;		
}

__global__ void paint(uchar4 *canvas, int *curPos, int *nextPos, const int *exitPos, const int DEBUGREP){
	uchar4 green, yellow, red;
	green.x = 0;	green.y = 255;	green.z = 0;	green.w = 255;
	yellow.x = 255;	yellow.y = 255;	yellow.z = 0;	yellow.w = 255;
	red.x = 255;	red.y = 0;		red.z = 0;		red.w = 0;
	
	int idx = threadIdx.x+blockIdx.x*blockDim.x;
	int paintRep = DEBUGREP;
	int code_pos = paintRep / REPNOPERLAYER * AGENTNO;
	if(idx<AGENTNO){
		int temp_agCurPos = curPos[code_pos + idx];
		int temp_agNextPos = nextPos[code_pos + idx];
		int temp_agExitPos = exitPos[code_pos + idx];
		int cpos = 0;
		int npos = 0;
		decode(cpos, temp_agCurPos, paintRep % REPNOPERLAYER);
		decode(npos, temp_agNextPos, paintRep % REPNOPERLAYER);
		if(cpos>=0){
			int cx = (int)cpos%DIM;
			int cy = (int)cpos/DIM;
			for(int i=0; i<CELLDIM-GAP; i++){
				for(int j=0; j<CELLDIM-GAP; j++){
					int offset = (CELLDIM*cx+i) + (CELLDIM*(DIM-1-cy)+j)*DIM*CELLDIM;
					canvas[offset]=green;
				}
			}
			cpos = exitPos[idx];
			decode(cpos, temp_agExitPos, paintRep % REPNOPERLAYER);
			cx = (int)cpos%DIM;
			cy = (int)cpos/DIM;
			for(int i=0; i<CELLDIM-GAP; i++){
				for(int j=0; j<CELLDIM-GAP; j++){
					int offset = (CELLDIM*cx+i) + (CELLDIM*(DIM-1-cy)+j)*DIM*CELLDIM;
					canvas[offset]=red;
				}
			}
		}
	}
}

__device__ bool combat(const int *env, const int agId, const int agNextPos, 
	const int *nextPos, const int *rands, const int repid){
		int rivalNextPos = 0;
		int rivalId = 0;
		int agNextX = agNextPos%DIM;
		int agNextY = agNextPos/DIM;
		if(agNextX-1 >= 0){
			decode(rivalId, env[agNextPos-1], repid);
			if(rivalId != EMPTY && rivalId != agId){
				decode(rivalNextPos, nextPos[rivalId], repid);
				int agRand = 0, rivalRand = 0;
				decode(agRand, rands[agId], repid);
				decode(rivalRand, rands[rivalId], repid);
				if(rivalNextPos == agNextPos && agRand <= rivalRand)
					return false;
			}
		}
		if(agNextX+1 < DIM){
			decode(rivalId, env[agNextPos+1], repid);
			if(rivalId != EMPTY && rivalId != agId){
				decode(rivalNextPos, nextPos[rivalId], repid);
				int agRand = 0, rivalRand = 0;
				decode(agRand, rands[agId], repid);
				decode(rivalRand, rands[rivalId], repid);
				if(rivalNextPos == agNextPos && agRand <= rivalRand)
					return false;
			}
		}
		if(agNextY-1 >= 0){
			decode(rivalId, env[agNextPos-DIM], repid);
			if(rivalId != EMPTY && rivalId != agId){
				decode(rivalNextPos, nextPos[rivalId], repid);
				int agRand = 0, rivalRand = 0;
				decode(agRand, rands[agId], repid);
				decode(rivalRand, rands[rivalId], repid);
				if(rivalNextPos == agNextPos && agRand <= rivalRand)
					return false;
			}
		}
		if(agNextY+1 < DIM){
			decode(rivalId, env[agNextPos+DIM], repid);
			if(rivalId != EMPTY && rivalId != agId){
				decode(rivalNextPos, nextPos[rivalId], repid);
				int agRand = 0, rivalRand = 0;
				decode(agRand, rands[agId], repid);
				decode(rivalRand, rands[rivalId], repid);
				if(rivalNextPos == agNextPos && agRand <= rivalRand)
					return false;
			}
		}
		return true;
}

__device__ int nextstep(const int *env, const int agExitPos, const int agCurPos, const int repid){
	int agNextPos = agCurPos;
	int agCurX = agCurPos%DIM;
	int agCurY = agCurPos/DIM;
	int agCurDistToExit = POS_TO_EXIT_2D(agCurX, agCurY, agExitPos);
	int agNextDistToExit = agCurDistToExit;
	int t = 0;
	if(agCurY-1 >= 0){
		int rep_envPos = 0;
		decode(rep_envPos, env[agCurPos-DIM], repid);
		if(rep_envPos == EMPTY){
			agNextDistToExit = POS_TO_EXIT_2D(agCurX, agCurY-1, agExitPos);
			t = TWO_TO_ONE(agCurX, agCurY-1);
			if(agNextDistToExit<agCurDistToExit){
				agNextPos = t;
				agCurDistToExit = agNextDistToExit;
			}
		}
	}
	if(agCurY+1 < DIM){
		int rep_envPos = 0;
		decode(rep_envPos, env[agCurPos+DIM], repid);
		if(rep_envPos == EMPTY){
			agNextDistToExit = POS_TO_EXIT_2D(agCurX, agCurY+1, agExitPos);
			t = TWO_TO_ONE(agCurX, agCurY+1);
			if(agNextDistToExit<agCurDistToExit){
				agNextPos = t;
				agCurDistToExit = agNextDistToExit;
			}
		}
	}
	if(agCurX-1 >= 0){
		int rep_envPos = 0;
		decode(rep_envPos, env[agCurPos-1], repid);
		if(rep_envPos == EMPTY){
			agNextDistToExit = POS_TO_EXIT_2D(agCurX-1, agCurY, agExitPos);
			t = TWO_TO_ONE(agCurX-1, agCurY);
			if(agNextDistToExit<agCurDistToExit){
				agNextPos = t;
				agCurDistToExit = agNextDistToExit;
			}
		}
	}
	if(agCurX+1 < DIM){
		int rep_envPos = 0;
		decode(rep_envPos, env[agCurPos+1], repid);
		if(rep_envPos == EMPTY){
			agNextDistToExit = POS_TO_EXIT_2D(agCurX+1, agCurY, agExitPos);
			t = TWO_TO_ONE(agCurX+1, agCurY);
			if(agNextDistToExit<agCurDistToExit){
				agNextPos = t;
				agCurDistToExit = agNextDistToExit;
			}
		}
	}
	return agNextPos;
}

extern __shared__ int smem[];
__global__ void inspection_kernel (const int *env, const int *curPos, int *nextPos, const int *rands, const bool *fini){
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	int acsidx = idx % AGENTNO; //access index, used to access array begins at code_pos
	if(acsidx<AGENTNO && idx < AGENTNO * LAYERNO){
		//int code_pos = repid / REPNOPERLAYER * AGENTNO;
		int code_pos = (int)idx/AGENTNO; code_pos *= AGENTNO;
		int rep_agCurPos = 0;
		int rep_agNextPos = 0;
		int rep_env = 0;
		//for(int repid=0; repid<REPNO && fini[repid] != true; repid++){
		for(int repid=0; repid<REPNOPERLAYER; repid++){
			decode(rep_agCurPos, curPos[code_pos + acsidx], repid % REPNOPERLAYER);
			decode(rep_agNextPos, nextPos[code_pos + acsidx], repid % REPNOPERLAYER);
			decode(rep_env, env[code_pos + rep_agNextPos], repid % REPNOPERLAYER);
			bool settledown = false;
			if (rep_env == EMPTY)
				settledown = combat(&env[code_pos], acsidx, rep_agNextPos, 
					&nextPos[code_pos], &rands[code_pos], repid % REPNOPERLAYER);
			if (settledown != true)
				rep_agNextPos = rep_agCurPos;
			encode(rep_agNextPos, nextPos[code_pos + acsidx], repid % REPNOPERLAYER);
		}
	}
}

__global__ void nextstep_kernel (const int *env, const int *exitPos, const int *curPos, int *nextPos, const bool *fini) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	int acsidx = idx % AGENTNO; //access index, used to access array begins at code_pos
	if(acsidx < AGENTNO && idx < AGENTNO * LAYERNO){
		//int code_pos = repid / REPNOPERLAYER * AGENTNO;
		int code_pos = (int)idx/AGENTNO; code_pos *= AGENTNO;
		int rep_agCurPos = 0;
		int rep_agExitPos = 0;
		//for(int repid=0; repid<REPNO && fini[repid] != true; repid++){
		for(int repid = 0; repid < REPNOPERLAYER; repid++){
			decode(rep_agCurPos, curPos[code_pos + acsidx], repid % REPNOPERLAYER);
			decode(rep_agExitPos, exitPos[code_pos + acsidx], repid % REPNOPERLAYER);
			int rep_agNextPos = nextstep(&env[code_pos], rep_agExitPos, rep_agCurPos, repid % REPNOPERLAYER);
			encode(rep_agNextPos, nextPos[code_pos + acsidx], repid % REPNOPERLAYER);
		}
	}
}

__global__ void register_kernel (int *env, const int *exitPos, int *curPos, const int *nextPos, bool *fini){
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	int acsidx = idx % AGENTNO; //access index, used to access array begins at code_pos
	if(acsidx < AGENTNO && idx < AGENTNO * LAYERNO){
		//int code_pos = repid / REPNOPERLAYER * AGENTNO;
		int code_pos = (int)idx/AGENTNO; code_pos *= AGENTNO;
		int rep_agCurPos = 0;
		int rep_agNextPos = 0;
		int rep_agExitPos = 0;
		int rep_env = 0;
		for(int repid = 0; repid < REPNOPERLAYER; repid++){
			decode(rep_agCurPos, curPos[code_pos + acsidx], repid % REPNOPERLAYER);
			decode(rep_agNextPos, nextPos[code_pos + acsidx], repid % REPNOPERLAYER);
			decode(rep_agExitPos, exitPos[code_pos + acsidx], repid % REPNOPERLAYER);
			rep_agCurPos = rep_agNextPos;
			rep_env = acsidx;
			if(rep_agCurPos == rep_agExitPos)
				rep_env = EMPTY; //when encoding -1 into integer, decoding it we obtain 1023.
			else
				fini[REPNOPERLAYER*idx/AGENTNO+repid] = false;
			encode(rep_agCurPos, curPos[code_pos + acsidx], repid % REPNOPERLAYER);
			encode(rep_env, env[code_pos + rep_agCurPos], repid % REPNOPERLAYER);
		}
	}
}

CUevent start, stop;
FILE	*fp;
void generate_frame( uchar4 *canvas, void*, GPUAnimBitmap &bitmap ) {
	timer_start(start, stop);
	dim3 washblock(DIM,DIM);
	dim3 washgrid(CELLDIM,CELLDIM);
	int blockSize = 128;
	int gridSize = LAYERNO * AGENTNO / blockSize + 1;

	cudaMemset(bitmap.env, -1, LAYERNO * DIM * DIM * sizeof(int));
	register_kernel<<<gridSize, blockSize>>>(bitmap.env, bitmap.exitPos, bitmap.curPos, bitmap.nextPos, bitmap.fini);
	nextstep_kernel<<<gridSize, blockSize>>>(bitmap.env, bitmap.exitPos, bitmap.curPos, bitmap.nextPos, bitmap.fini);
	inspection_kernel<<<gridSize, blockSize>>>(bitmap.env, bitmap.curPos, bitmap.nextPos, (int *)bitmap.rands, bitmap.fini);
	wash<<<washgrid, washblock>>>(canvas);
	paint<<<gridSize, blockSize>>>(canvas, bitmap.curPos, bitmap.nextPos, bitmap.exitPos, DEBUGREP);

	float oneIterTime = timer_stop(start, stop);
	cleanup<<<1, blockSize>>>(bitmap.fini, bitmap.runningLoop, bitmap.runningTime, oneIterTime);
	
	//fp = fopen("log.txt", "a");
	//int *hdata = (int*)malloc(AGENTNO*sizeof(int));
	//cudaMemcpy(hdata, &bitmap.env[DEBUGREP/REPNOPERLAYER*AGENTNO], AGENTNO*sizeof(int), cudaMemcpyDeviceToHost);
	//fprintf(fp, "env\n");
	//for(int i=0; i<AGENTNO; i++){
	//	int pos = 0;
	//	decode_host(pos, hdata[i], DEBUGREP % REPNOPERLAYER);
	//	fprintf(fp, "%4d ", pos);
	//	if(i%(DIM)==DIM-1)
	//		fprintf(fp,"\n");
	//	fflush(fp);
	//}
	//cudaMemcpy(hdata, &bitmap.nextPos[DEBUGREP/REPNOPERLAYER*AGENTNO], AGENTNO*sizeof(int), cudaMemcpyDeviceToHost);
	//fprintf(fp, "nextPos\n");
	//for(int i=0; i<AGENTNO; i++){
	//	int pos = 0;
	//	decode_host(pos, hdata[i], DEBUGREP % REPNOPERLAYER);
	//	fprintf(fp, "%4d ", pos);
	//	if(i%(DIM)==DIM-1)
	//		fprintf(fp,"\n");
	//	fflush(fp);
	//}
	//cudaMemcpy(hdata, &bitmap.rands[DEBUGREP/REPNOPERLAYER*AGENTNO], AGENTNO*sizeof(int), cudaMemcpyDeviceToHost);
	//fprintf(fp, "rands\n");
	//for(int i=0; i<AGENTNO; i++){
	//	int pos = 0;
	//	decode_host(pos, hdata[i], DEBUGREP % REPNOPERLAYER);
	//	fprintf(fp, "%4d ", pos);
	//	if(i%(DIM)==DIM-1)
	//		fprintf(fp,"\n");
	//	fflush(fp);
	//}
	//fclose(fp);
	//free(hdata);
}

__global__ void init(int* nextPos, int *rands){
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx<AGENTNO)
		nextPos[idx]=(int)(rands[idx]*DIM*DIM);
}

__global__ void init2(int* nextPos, unsigned int *rands){
	const int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx<AGENTNO){
		for(int i = 0; i < REPNO; i++){
			int code_pos = i / REPNOPERLAYER * AGENTNO;
			encode(idx, nextPos[code_pos + idx], i % REPNOPERLAYER);
		}
	}
}

__global__ void initexit(const int *nextPos, int *exitPos, const int *exitList, const int exitAmount){
	const int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx<AGENTNO){
		int rep_agNextPos = 0;
		int rep_agExitPos = 0;
		int code_pos = 0;
		for(int i=0; i<REPNO; i++){ // iterate all reps
			code_pos = i / REPNOPERLAYER * AGENTNO;
			decode(rep_agNextPos, nextPos[code_pos + idx], i % REPNOPERLAYER);
			int dist = POS_TO_EXIT_1D(rep_agNextPos, 0);
			rep_agExitPos = exitList[0];
			for(int j = 1; j < exitAmount; j++){ // iterate exitList
				int dist2 = POS_TO_EXIT_1D(rep_agNextPos, exitList[j]);
				if(dist > dist2 && j <= i){
					dist = dist2;
					rep_agExitPos = exitList[j];
				}
			}
			encode(rep_agExitPos, exitPos[code_pos + idx], i % REPNOPERLAYER);
		}
	}
}

int *hexit = new int[100];
void init_state(GPUAnimBitmap &bitmap){
	int blocksize = 32;
	int gridsize = AGENTNO/blocksize+1;
	cudaMalloc((void**)&(bitmap.rands), LAYERNO * AGENTNO * sizeof(float));
	cudaMalloc((void**)&(bitmap.curPos), LAYERNO * AGENTNO * sizeof(int));
	cudaMalloc((void**)&(bitmap.nextPos), LAYERNO * AGENTNO * sizeof(int));
	cudaMalloc((void**)&(bitmap.exitPos), LAYERNO * AGENTNO * sizeof(int));
	cudaMalloc((void**)&(bitmap.env), LAYERNO * DIM * DIM * sizeof(int));
	cudaMalloc((void**)&(bitmap.fini), REPNO * sizeof(bool));
	cudaMalloc((void**)&(bitmap.runningLoop), REPNO * sizeof(int));
	cudaMalloc((void**)&(bitmap.runningTime), REPNO * sizeof(float));
	curandCreateGenerator(&(bitmap.gen), CURAND_RNG_PSEUDO_DEFAULT);
	curandSetPseudoRandomGeneratorSeed(bitmap.gen, 2345);
	curandGenerate(bitmap.gen, bitmap.rands, LAYERNO * AGENTNO);
	cudaMemset(bitmap.env, EMPTY, LAYERNO * DIM * DIM * sizeof(int));
	cudaMemset(bitmap.fini, true, REPNO * sizeof(bool));
	cudaMemset(bitmap.runningTime, 0, REPNO * sizeof(float));
	cudaMemset(bitmap.runningLoop, 0, REPNO * sizeof(int));
	init2<<<gridsize,blocksize>>>(bitmap.nextPos ,bitmap.rands);
	
	//int *hdata = (int*)malloc(AGENTNO*sizeof(int));
	//int *hdata2 = (int*)malloc(DIM*DIM*sizeof(int));
	//memset(hdata2, 0, DIM*DIM*sizeof(int));
	//cudaMemcpy(hdata, &bitmap.nextPos[AGENTNO], AGENTNO*sizeof(int), cudaMemcpyDeviceToHost);
	//for(int i=0; i<AGENTNO; i++){
	//	int pos = 0;
	//	decode_host(pos, hdata[i], DEBUGREP % REPNOPERLAYER);
	//	printf("%4d ", pos);
	//	if(i%DIM==DIM-1)
	//		printf("\n");
	//}
	//for(int i=0; i<AGENTNO; i++){
	//	if(hdata2[hdata[i]] == 0)
	//		hdata2[hdata[i]] = 1;
	//	else {
	//		int j = hdata[i];
	//		while(hdata2[j] != 0){
	//			j++;
	//			if(j>=DIM*DIM)
	//				j=0;
	//		}
	//		hdata2[j]=1;
	//		hdata[i]=j;
	//	}
	//	int j = hdata[i];
	//	//printf("pos:%d, pos_x:%d, pos_y:%d, dist:%d\n", hdata[i], hdata[i]%DIM, hdata[i]/DIM, TWO_TO_EXIT(hdata[i]%DIM,hdata[i]/DIM));
	//}
	//cudaMemcpy(bitmap.state2, hdata, AGENTNO*sizeof(int), cudaMemcpyHostToDevice);
	int *dexit;
	memset(hexit, 0, 100*sizeof(int));
	hexit[0] = TWO_TO_ONE(0, 0);
	if(EXITNO>1)
		hexit[1] = TWO_TO_ONE(0, DIM-1);
	if(EXITNO>2)
		hexit[2] = TWO_TO_ONE(DIM-1, 0);
	if(EXITNO>3)
		hexit[3] = TWO_TO_ONE(DIM-1, DIM-1);
	if(EXITNO>4)
		hexit[4] = TWO_TO_ONE(DIM/2, DIM/2);
	if(EXITNO>5)
		hexit[5] = TWO_TO_ONE(DIM/3, DIM/3);
	if(EXITNO>6)
		hexit[6] = TWO_TO_ONE(2*DIM/3, 2*DIM/3);
	if(EXITNO>7)
		hexit[7] = TWO_TO_ONE(DIM/3, 2*DIM/3);
	if(EXITNO>8)
		hexit[8] = TWO_TO_ONE(2*DIM/3, DIM/3);
	if(EXITNO>9)
		hexit[9] = TWO_TO_ONE(DIM/4, DIM/4);
	//for(int i=0; i<EXITNO; i++)
	//	printf("%4d ", hexit[i]);
	//printf("\n");
	cudaMalloc((void**)&dexit, EXITNO*sizeof(int));
	cudaMemcpy(dexit, hexit, EXITNO*sizeof(int), cudaMemcpyHostToDevice);
	initexit<<<gridsize, blocksize>>>(bitmap.nextPos, bitmap.exitPos, dexit, EXITNO);
	//int *hdata = (int*)malloc(AGENTNO*sizeof(int));
	//cudaMemcpy(hdata, &bitmap.exitPos[AGENTNO], AGENTNO*sizeof(int), cudaMemcpyDeviceToHost);
	//for(int i=0; i<AGENTNO; i++){
	//	int pos = 0;
	//	decode_host(pos, hdata[i], 2);
	//	printf("%4d ", pos);
	//	if(i%DIM==DIM-1)
	//		printf("\n");
	//}
	//free(hdata);
	//free(hdata2);
	cudaFree(dexit);	
}

void decision_point(int x, int y, int *nextPos, int *exitPos){
	printf("%d, %d -> ", x, y);
	int blocksize = 32;
	int gridsize = AGENTNO/blocksize+1;
	int celldim = CELLDIM;
	x = x/celldim;
	y = y/celldim;
	printf("%d, %d\n", x, y);

	hexit[EXITNO] = TWO_TO_ONE(x, y);
	int *dexit;
	cudaMalloc((void**)&dexit, (EXITNO+1)*sizeof(int));
	EXITNO++;
	cudaMemcpy(dexit, hexit, EXITNO*sizeof(int), cudaMemcpyHostToDevice);
	initexit<<<gridsize, blocksize>>>(nextPos, exitPos, dexit, EXITNO);
}
/*compact the clones in the space*/